__device__ void
BWT_Encode(unsigned int /*[in]*/bytes, const unsigned char * /*[in]*/data, unsigned char * /*[out]*/encoded, unsigned int & /*[out]*/dec_start_pos)
{
    unsigned char *buffer = new unsigned char[2*Config::BLOCK_SIZE+1];       // alloc twice as much because I need to duplicate the buffer for comparison; then add 1 extra byte for the last byte (predecessor of the first byte on a block)
    unsigned char **ptrs = new unsigned char*[Config::BLOCK_SIZE], **tmp = new unsigned char*[Config::BLOCK_SIZE];
    // encode a single block
    if(bytes > 0)
    {
        memcpy(buffer + 1, data, bytes);
        memcpy(buffer + bytes + 1, buffer + 1, bytes);    // duplicate the buffer
        buffer[0] = buffer[bytes];       // add 1 extra byte for the last byte (predecessor of the first byte on a block)
        for(streamsize i = 0; i < bytes; i++)
            ptrs[i] = buffer + i + 1;
//        mergeSort(bytes, ptrs, tmp, 0, bytes-1, compare_str);
        for(streamsize i = 0; i < bytes; i++)
        {
            encoded[i] = *(ptrs[i] - 1);   // all that extra byte stuff is because of this ;-)
            if(ptrs[i] == buffer+1) dec_start_pos = i;
        }
    }
    delete [] buffer;
    delete [] ptrs;
    delete [] tmp;
}

__device__ void
BWT_Decode(unsigned int /*[in]*/bytes, const unsigned char * /*[in]*/data, unsigned int /*[in]*/dec_start_pos, unsigned char * /*[out]*/decoded)
{
    unsigned char **ptrs = new unsigned char*[Config::BLOCK_SIZE], **tmp = new unsigned char*[Config::BLOCK_SIZE];
    unsigned int *LF = new unsigned int[Config::BLOCK_SIZE];
    if(bytes > 0)
    {
        // compute LF
        for(streamsize i = 0; i < bytes; i++)
            ptrs[i] = (unsigned char *)(data + i);
//        mergeSort(bytes, ptrs, tmp, 0, bytes-1, compare_char);
        for(streamsize i = 0; i < bytes; i++)
            LF[ptrs[i]-data] = i;
        // decode
        streamsize dec_i = bytes - 1;
        decoded[dec_i] = data[dec_start_pos];
        while(dec_i > 0)
            decoded[--dec_i] = data[dec_start_pos=LF[dec_start_pos]];
    }
    delete [] ptrs;
    delete [] tmp;
    delete [] LF;
}


__device__ bool
compare_str(int n, unsigned char *a, unsigned char *b)
{
    int i = 0;
    while((a[i] == b[i]) && (i < n)) ++i;
    return (a[i] <= b[i]);
}

__device__ bool
compare_char(int n, unsigned char *a, unsigned char *b)
{
    return ((*a) <= (*b));
}

/**
 * Slevani pro Merge sort
 * @param array pole k serazeni
 * @param aux pomocne pole (stejne velikosti jako razene)
 * @param left prvni index, na ktery smim sahnout
 * @param right posledni index, na ktery smim sahnout
 */
__device__ void
merge(int n, unsigned char *array[], unsigned char *aux[], int left, int right, bool (*cmp)(int, unsigned char *, unsigned char *)) {
    int middleIndex = (left + right)/2;
    int leftIndex = left; 
    int rightIndex = middleIndex + 1;
    int auxIndex = left;
    while (leftIndex <= middleIndex && rightIndex <= right) {
        if (cmp(n, array[leftIndex], array[rightIndex])) {
            aux[auxIndex] = array[leftIndex++];
        } else {
            aux[auxIndex] = array[rightIndex++];
        }
        auxIndex++;
    }
    while (leftIndex <= middleIndex) {
        aux[auxIndex] = array[leftIndex++];
        auxIndex++;
    }
    while (rightIndex <= right) {
        aux[auxIndex] = array[rightIndex++];
        auxIndex++;
    }
}    

/** 
 * Razeni slevanim (od nejvyssiho)
 * @param array pole k serazeni
 * @param aux pomocne pole stejne delky jako array
 * @param left prvni index na ktery se smi sahnout
 * @param right posledni index, na ktery se smi sahnout
 */
__device__ void
mergeSort(int n, unsigned char *array[], unsigned char *aux[], int left, int right, bool (*cmp)(int, unsigned char *, unsigned char *)) {
    if (left == right) return; 
    int middleIndex = (left + right)/2;
    mergeSort(n, array, aux, left, middleIndex, cmp);
    mergeSort(n, array, aux, middleIndex + 1, right, cmp);
    merge(n, array, aux, left, right, cmp);
  
    for (int i = left; i <= right; i++) {
        array[i] = aux[i];
    }
}